#include <cutil_inline.h>
#include <device_launch_parameters.h>
#include <GL\freeglut.h>
#include <thrust\sort.h>
#include <thrust\device_vector.h>
#include <thrust\functional.h>
#include <thrust\transform.h>
#include "config.cuh"

typedef int test_t;

__global__ void init_kernel(int* d_index, int* d_init, int* d_baseId, int* d_no, int* d_etype, 
	float* d_time, float* d_duration, float* d_speed, float* d_pos, int N)
{
	unsigned int i = threadIdx.x+blockIdx.x*blockDim.x;
	if(i<N){
		int eidx = d_init[i];
		int no = d_no[eidx];
		float time = d_time[eidx];
		float duration = d_duration[eidx];
		float speed = d_speed[eidx];
		float pos = d_pos[eidx];
		//d_time[eidx]= time;
		d_index[eidx]=eidx;

		float termiPos = pos+duration*speed/3600;
		if(termiPos>40)
			termiPos = 40;
		float currPos = pos;
		float nextPos = ((int)(currPos/2)+1)*2;
		while(nextPos<termiPos){
			eidx++;
			d_index[eidx] = eidx;
			d_time[eidx] = time = time+(nextPos-currPos)/speed*3600;
			d_speed[eidx] = speed;
			d_pos[eidx] = nextPos;
			d_no[eidx] = no;
			d_duration[eidx] = 0;
			d_etype[eidx] = TERMI;
			d_baseId[eidx] = currPos/2;

			eidx++;
			d_index[eidx] = eidx;
			d_time[eidx] = time;
			d_speed[eidx] = speed;
			d_pos[eidx] = nextPos;
			d_no[eidx] = no;
			d_duration[eidx] = duration-(nextPos-pos)/speed*3600;
			d_etype[eidx] = HANDO;
			d_baseId[eidx] = nextPos/2;

			currPos = nextPos;
			nextPos = nextPos+2;
		}
		eidx++;
		d_index[eidx] = eidx;
		d_time[eidx] = time+(termiPos-currPos)/speed*3600;
		d_speed[eidx] = speed;
		d_pos[eidx] = termiPos;
		d_no[eidx] = no;
		d_duration[eidx] = 0;
		d_etype[eidx] = TERMI;
		d_baseId[eidx] = currPos/2;
	}
}

//float* d_init, float* d_duration, float* d_speed, float* d_pos can be freed
extern "C" void launch_init_kernel(da_t da, int n)
{
	//n = SAMPLENO;
	init_kernel<<<1+n/BLOCKSIZE,BLOCKSIZE>>>(da.d_index, da.d_init, da.d_baseId, da.d_no, da.d_etype, 
		da.d_time, da.d_duration, da.d_speed, da.d_pos, n);
	cudaFree(da.d_speed);
	cudaFree(da.d_pos);
	cudaFree(da.d_duration);
	cudaFree(da.d_init);
}

typedef thrust::tuple<int, float, int> KeyType;
struct Comp : public thrust::binary_function<KeyType, KeyType, bool>
{
	__host__ __device__
		bool operator () (const KeyType& a, const KeyType& b) const
	{
		if(thrust::get<0>(a)!=thrust::get<0>(b))
			return thrust::get<0>(a) < thrust::get<0>(b);
		else if(thrust::get<1>(a)!=thrust::get<1>(b))
			return thrust::get<1>(a) < thrust::get<1>(b);
		else
			return thrust::get<2>(a) < thrust::get<2>(b);
	}
};

//float *d_time can be freed
extern "C" void launch_sort_kernel(da_t &da, int n)
{
	thrust::device_ptr<int> d_etype_ptr(da.d_etype);
	thrust::device_ptr<int> d_baseId_ptr(da.d_baseId);
	thrust::device_ptr<int> d_no_ptr(da.d_no);
	thrust::device_ptr<float> d_time_ptr(da.d_time);

	typedef thrust::device_vector<int>::iterator IntIter;
	typedef thrust::device_vector<float>::iterator FloatIter;
	typedef thrust::tuple<IntIter, FloatIter, IntIter> KeyIterTuple; //base, time, no
	typedef thrust::tuple<IntIter> ValIterTuple; //etype
	typedef thrust::zip_iterator<KeyIterTuple> KeyIter;
	typedef thrust::zip_iterator<ValIterTuple> ValIter;

	KeyIter key_begin(thrust::make_tuple(d_baseId_ptr, d_time_ptr, d_no_ptr));
	KeyIter key_end(thrust::make_tuple(d_baseId_ptr+n, d_time_ptr+n, d_no_ptr+n));
	ValIter val_begin(thrust::make_tuple(d_etype_ptr));

	thrust::sort_by_key(key_begin, key_end, val_begin, Comp());
	cudaFree(da.d_time);
}

extern __shared__ int smem[];
__global__ void correction_kernel(int* d_etype, int n){
	const unsigned int tid	= threadIdx.x;
	const unsigned int i	= threadIdx.x+blockDim.x*blockIdx.x;
	int	* s_etype = sem;
	if(i<n)
		s_etype[tid]=d_etype[i];

}

extern "C" void launch_correction_kernel(da_t &da, int n){
}